#include <stdio.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>

static __global__ void testKernel(float* data) {
    // printf("===threadIdx.x = %d\n", threadIdx.x);
    printf("===threadIdx.x = %d, data[%d] = %f\n", threadIdx.x, threadIdx.x, data[threadIdx.x]);
}


// cudaMallocManaged 是 CUDA 统一内存管理(Unified Memory)的核心接口，它简化了 CPU 和 GPU 之间的内存管理。
// 1.**统一内存空间**：
// - 单个指针同时在CPU和GPU上有效
// - 无需显式的cudaMemcpy操作
// 2.**按需迁移**：
// - 数据在CPU或GPU首次访问时自动迁移
// - 支持页面级粒度迁移 (从Pascal架构开始)
// 3.**一致性模型**：
// - 默认弱一致性 (需要同步点)
// - 可使用cudaDeviceSynchronize()强制同步

int main() {
    // 分配100个float的托管内存
    float* data;
    cudaMallocManaged(&data, 100 * sizeof(float));

    // CPU端初始化
    for(int i = 0; i < 100; i++) {
        data[i] = i;
    }

    // GPU端处理
    testKernel<<<1, 100>>>(data);

    // host等待device核函数执行完成
    cudaDeviceSynchronize();

    // CPU端可直接访问结果
    printf("data[0] = %f\n", data[0]);

    // 释放内存
    cudaFree(data);


    return 0;
}


// 可以看到如下信息，0~31连续打印，32~63连续打印，64~95连续打印，96~99连续打印。因为gpu执行核函数时，
// 每个warp是执行的最小单元，每个warp包含32个线程，SIMT保证，在同一个warp中，保证每个线程执行相同的指令。
// 所以，每32行（即同一个warp）连续打印。不足32整数倍，也会分配一个warp来执行。

// ===threadIdx.x = 96, data[96] = 96.000000
// ===threadIdx.x = 97, data[97] = 97.000000
// ===threadIdx.x = 98, data[98] = 98.000000
// ===threadIdx.x = 99, data[99] = 99.000000
// ===threadIdx.x = 32, data[32] = 32.000000
// ===threadIdx.x = 33, data[33] = 33.000000
// ===threadIdx.x = 34, data[34] = 34.000000
// ===threadIdx.x = 35, data[35] = 35.000000
// ===threadIdx.x = 36, data[36] = 36.000000
// ===threadIdx.x = 37, data[37] = 37.000000
// ===threadIdx.x = 38, data[38] = 38.000000
// ===threadIdx.x = 39, data[39] = 39.000000
// ===threadIdx.x = 40, data[40] = 40.000000
// ===threadIdx.x = 41, data[41] = 41.000000
// ===threadIdx.x = 42, data[42] = 42.000000
// ===threadIdx.x = 43, data[43] = 43.000000
// ===threadIdx.x = 44, data[44] = 44.000000
// ===threadIdx.x = 45, data[45] = 45.000000
// ===threadIdx.x = 46, data[46] = 46.000000
// ===threadIdx.x = 47, data[47] = 47.000000
// ===threadIdx.x = 48, data[48] = 48.000000
// ===threadIdx.x = 49, data[49] = 49.000000
// ===threadIdx.x = 50, data[50] = 50.000000
// ===threadIdx.x = 51, data[51] = 51.000000
// ===threadIdx.x = 52, data[52] = 52.000000
// ===threadIdx.x = 53, data[53] = 53.000000
// ===threadIdx.x = 54, data[54] = 54.000000
// ===threadIdx.x = 55, data[55] = 55.000000
// ===threadIdx.x = 56, data[56] = 56.000000
// ===threadIdx.x = 57, data[57] = 57.000000
// ===threadIdx.x = 58, data[58] = 58.000000
// ===threadIdx.x = 59, data[59] = 59.000000
// ===threadIdx.x = 60, data[60] = 60.000000
// ===threadIdx.x = 61, data[61] = 61.000000
// ===threadIdx.x = 62, data[62] = 62.000000
// ===threadIdx.x = 63, data[63] = 63.000000
// ===threadIdx.x = 0, data[0] = 0.000000
// ===threadIdx.x = 1, data[1] = 1.000000
// ===threadIdx.x = 2, data[2] = 2.000000
// ===threadIdx.x = 3, data[3] = 3.000000
// ===threadIdx.x = 4, data[4] = 4.000000
// ===threadIdx.x = 5, data[5] = 5.000000
// ===threadIdx.x = 6, data[6] = 6.000000
// ===threadIdx.x = 7, data[7] = 7.000000
// ===threadIdx.x = 8, data[8] = 8.000000
// ===threadIdx.x = 9, data[9] = 9.000000
// ===threadIdx.x = 10, data[10] = 10.000000
// ===threadIdx.x = 11, data[11] = 11.000000
// ===threadIdx.x = 12, data[12] = 12.000000
// ===threadIdx.x = 13, data[13] = 13.000000
// ===threadIdx.x = 14, data[14] = 14.000000
// ===threadIdx.x = 15, data[15] = 15.000000
// ===threadIdx.x = 16, data[16] = 16.000000
// ===threadIdx.x = 17, data[17] = 17.000000
// ===threadIdx.x = 18, data[18] = 18.000000
// ===threadIdx.x = 19, data[19] = 19.000000
// ===threadIdx.x = 20, data[20] = 20.000000
// ===threadIdx.x = 21, data[21] = 21.000000
// ===threadIdx.x = 22, data[22] = 22.000000
// ===threadIdx.x = 23, data[23] = 23.000000
// ===threadIdx.x = 24, data[24] = 24.000000
// ===threadIdx.x = 25, data[25] = 25.000000
// ===threadIdx.x = 26, data[26] = 26.000000
// ===threadIdx.x = 27, data[27] = 27.000000
// ===threadIdx.x = 28, data[28] = 28.000000
// ===threadIdx.x = 29, data[29] = 29.000000
// ===threadIdx.x = 30, data[30] = 30.000000
// ===threadIdx.x = 31, data[31] = 31.000000
// ===threadIdx.x = 64, data[64] = 64.000000
// ===threadIdx.x = 65, data[65] = 65.000000
// ===threadIdx.x = 66, data[66] = 66.000000
// ===threadIdx.x = 67, data[67] = 67.000000
// ===threadIdx.x = 68, data[68] = 68.000000
// ===threadIdx.x = 69, data[69] = 69.000000
// ===threadIdx.x = 70, data[70] = 70.000000
// ===threadIdx.x = 71, data[71] = 71.000000
// ===threadIdx.x = 72, data[72] = 72.000000
// ===threadIdx.x = 73, data[73] = 73.000000
// ===threadIdx.x = 74, data[74] = 74.000000
// ===threadIdx.x = 75, data[75] = 75.000000
// ===threadIdx.x = 76, data[76] = 76.000000
// ===threadIdx.x = 77, data[77] = 77.000000
// ===threadIdx.x = 78, data[78] = 78.000000
// ===threadIdx.x = 79, data[79] = 79.000000
// ===threadIdx.x = 80, data[80] = 80.000000
// ===threadIdx.x = 81, data[81] = 81.000000
// ===threadIdx.x = 82, data[82] = 82.000000
// ===threadIdx.x = 83, data[83] = 83.000000
// ===threadIdx.x = 84, data[84] = 84.000000
// ===threadIdx.x = 85, data[85] = 85.000000
// ===threadIdx.x = 86, data[86] = 86.000000
// ===threadIdx.x = 87, data[87] = 87.000000
// ===threadIdx.x = 88, data[88] = 88.000000
// ===threadIdx.x = 89, data[89] = 89.000000
// ===threadIdx.x = 90, data[90] = 90.000000
// ===threadIdx.x = 91, data[91] = 91.000000
// ===threadIdx.x = 92, data[92] = 92.000000
// ===threadIdx.x = 93, data[93] = 93.000000
// ===threadIdx.x = 94, data[94] = 94.000000
// ===threadIdx.x = 95, data[95] = 95.000000
// data[0] = 0.000000